/*
 * Copyright (c) 2010, GPU group, BTLab, SJTU.
 * All rights reserved.
 *
 * Filename: PTXTranslator.cpp
 * Abstract: Complete definition of classes GVInst and PTXTranslator.
 *
 * Current version: 1.0
 * Programmer: Dylan Dong
 * Data: 2010-01-24
 *
 * Replace version: 1.1
 * Programmer: Dylan Dong
 * Data: 2010-03-12
 * PS: Add supportion of FFT
 */

#include "g_crossbit/PTXTranslator.h"

using namespace g_crossbit;

#define FLOAT
/*
 * Function Introduce: Used for setting gvinst's m_op. 
 * 			  		   The value setted is according to the string input.
 *
 * Input Parameter: string
 * Returned Value: Opcode enum-type to describe gvinst's opcode
 */
Opcode SetOpcode(const string &op_str)
{
	Opcode op;
	//op_str = op_str.substr(0, op_str.find('.'));
	if(op_str == "GET")
		op = OP_GET;
	else if (op_str == "PUT")
		op = OP_PUT;
	else if (op_str == "MOV")
		op = OP_MOV;
	else if (op_str == "LD")
		op = OP_LD;
	else if (op_str == "ST")
		op = OP_ST;
	else if (op_str == "LI")
		op = OP_LI;
	else if (op_str == "ADD")
		op = OP_ADD;
	else if (op_str == "ADDI")
		op = OP_ADDI;
	else if (op_str == "SUB")
		op = OP_SUB;
	else if (op_str == "SUBI")
		op = OP_SUBI;
	else if (op_str == "MUL")
		op = OP_MUL;
	else if (op_str == "SLL")
		op = OP_SLL;
	else if (op_str == "AND")
		op = OP_AND;
	/*
	else if (op_str == "GETD")
		op = OP_GETD;
	else if (op_str == "PUTD")
		op = OP_PUTD;
	*/
	else if (op_str == "FLD")
		op = OP_FLD;
	else if (op_str == "FST")
		op = OP_FST;
	else if (op_str == "CONV")
		op = OP_CONV;
	else if (op_str == "CMP")
		op = OP_CMP;
	else if (op_str == "BRANCH")
		op = OP_BRANCH;
	else if (op_str == "entry:")
		op = OP_ENTRY;
	else if (op_str == "next:")
		op = OP_NEXT;
	else 
		op = OP_ERR; 
	return op;
}

/*
 * Function Introduce: Used for setting gvinst's m_con_code.
 * 			  		   The value setted is according to the string input.
 *
 * Input Parameters: string
 * Returned Value: CC enum-type to describe gvinst's conditio code
 */
CC SetCC(const string &cc_str)
{
	CC cc;
	if(cc_str == "EQ")
		cc = eq;
	else if (cc_str == "NE")
		cc = ne;
	else if (cc_str == "LT")
		cc = lt_s32;
	else if (cc_str == "BE")
		cc = lt_u32;
	else if (cc_str == "LE")
		cc = le;
	else
		cc = noncc;
	return cc;
}

/*
 * Function Introduce: Used for setting gvinst's m_stype and m_dtype.
 * 			  		   The value setted is according to the string input.
 *
 * Input Parameters: string
 * Returned Value: CC enum-type to describe gvinst's conditio code
 */
DataType SetDataType(const string &dt_str)
{
	DataType dt;
	if(dt_str == "u32")
		dt = u32;
	else if (dt_str == "u64")
		dt = u64;
	else if (dt_str == "s32")
		dt = s32;
	else if (dt_str == "s64")
		dt = s64;	
	else if (dt_str == "f32")
		dt = f32;
	else if (dt_str == "f64")
		dt = f64;
	else
		dt = notype;
	return dt;
}

/* 
 * static data member of class GVInst 
 * Record the number of class GVInst's object.
 */
int GVInst::s_gvinst_num = 0;

/*
 * Function Introduce: class GVInst's constructor.
 * 
 * Input Parameters: list<string>, containing a GVInst instruction.
 */
GVInst::GVInst(const list<string> &line_of_gvinst)
{
	list<string>::const_iterator gvinst_it = line_of_gvinst.begin();
	list<string>::const_reverse_iterator r_gvinst_it = line_of_gvinst.rbegin();
	
	m_op = SetOpcode(*gvinst_it);
	++gvinst_it;

	switch(m_op)
	{
		case OP_GET:
		case OP_PUT:
		case OP_MOV:
		case OP_LD:
		case OP_ST:
		case OP_LI:
			m_stype = m_dtype = SetDataType(*gvinst_it);
			m_soperand1 = *(++gvinst_it);
			m_soperand2 = "NULL";
			m_doperand = *r_gvinst_it;
			break;
		case OP_ADD:
		case OP_ADDI:
		case OP_SUB:
		case OP_SUBI:
		case OP_MUL:
		case OP_SLL:
		case OP_AND:
			m_stype = m_dtype = SetDataType(*gvinst_it);
			m_soperand1 = *(++gvinst_it);
			m_soperand2 = *(++gvinst_it);
			m_doperand = *r_gvinst_it;
			break;
		case OP_CONV:
			m_stype = SetDataType(*gvinst_it);
			m_dtype = SetDataType(*(++gvinst_it));
			m_soperand1 = *(++gvinst_it);
			m_soperand2 = *(++gvinst_it);
			m_doperand = *r_gvinst_it;
			break;		
		case OP_CMP:
		case OP_BRANCH:
			m_dtype = s32;	
			m_con_code = SetCC(*gvinst_it);
			m_soperand1 = *(++gvinst_it);
			m_soperand2 = *(++gvinst_it);
			m_doperand = *r_gvinst_it;
			break;
		case OP_ENTRY:
		case OP_NEXT:
			m_doperand = *r_gvinst_it;
			break;
		default:
			break;
	}
	
	GVInst::IncGVInstNum();
}

/*
 * Function Introduce: Allocate GPU's registers to gvinst's operands.
 * 
 * Input Parameters: GVInst gvinst
 * Output Parameters: string dreg, string sreg
 * 		dreg is one of GPU's registers which is allocated to gvinst's destination operand. 
 * 		sreg corresponds to source opedrand. 
 * Returned Value: return_flag
 * 		If any gvisnt's operand show a address in base addressing mode, like "v29(0x0)", then return 0,
 *		else, return 1;
 */
int PTXTranslator::SetRegs(const GVInst &gvinst, string &dreg, string &sreg)
{
	string reg, base;
	int count, pos, return_flag = 1;
	map<string, string> reg_table;
	map<string, string>::iterator map_it;

	// Set sreg
	switch (gvinst.GetSType())
	{
		case s32:
			if (gvinst.GetSoperand1().substr(0, 4) == "TEMP")
			{
				reg = "%f";
				count = m_f_count;
				reg_table = m_reg_table;
				return_flag = 3;
				break;
			}
			reg = "%r";
			count = m_r_count;
			reg_table = m_reg_table;
			break;
		case f32:
			reg = "%f";
			count = m_f_count;
			reg_table = m_reg_table;
			break;
		case f64:
			reg = "%fd";
			count = m_fd_count;
			reg_table = m_reg_table_64;
			break;
		default:
			break;
	}
	if (gvinst.GetSoperand1().substr(0, 2) == "0x")
		sreg = gvinst.GetSoperand1();
	else if (gvinst.GetSoperand1() == "v0")
		sreg = "0x0";
	else if (gvinst.GetSoperand1() == "f0")
		sreg = "0f00000000";
	else 
	{
		base = gvinst.GetSoperand1();
		pos = base.find('(');
		if (pos != string::npos)
		{
			base = base.substr(0, pos);
			return_flag = 0;
		}
		
		map_it = reg_table.find(base);
		if (map_it != reg_table.end())
			sreg = map_it->second;
		else
		{
			sreg = reg + int2str(count);
			reg_table.insert(make_pair(base, sreg));
			++count;
		}
	}

	switch (gvinst.GetSType())
	{
		case s32:
			if (return_flag == 3)
			{
				m_f_count = count;
				m_reg_table = reg_table;
				break;
			}
			m_r_count = count;
			m_reg_table = reg_table;
			break;
		case f32:
			m_f_count = count;
			m_reg_table = reg_table;
			break;
		case f64:
			m_fd_count = count;
			m_reg_table_64 = reg_table;
			break;
		default:
			break;
	}
	// Set dreg
	switch (gvinst.GetDType())
	{
		case s32:
			if (gvinst.GetDoperand().substr(0, 4) == "TEMP")
			{
				reg = "%f";
				count = m_f_count;
				reg_table = m_reg_table;
				return_flag = 2;
				break;
			}
			reg = "%r";
			count = m_r_count;
			reg_table = m_reg_table;
			break;
		case f32:
			reg = "%f";
			count = m_f_count;
			reg_table = m_reg_table;
			break;
		case f64:
			reg = "%fd";
			count = m_fd_count;
			reg_table = m_reg_table_64;
			break;
		default:
			break;
	}
	base = gvinst.GetDoperand();
	pos = base.find('(');
	if (pos != string::npos)
	{
		base = base.substr(0, pos);
		return_flag = 0;
	}
	map_it = reg_table.find(base);
	if (map_it != reg_table.end())
	{
		dreg = map_it->second;
	}
	else
	{
		dreg = reg + int2str(count);
		reg_table.insert(make_pair(base, dreg));
		++count;
	}

	switch (gvinst.GetDType())
	{
		case s32:
			if (return_flag == 2)
			{
				m_f_count = count;
				m_reg_table = reg_table;
				break;
			}
			m_r_count = count;
			m_reg_table = reg_table;
			break;
		case f32:
			m_f_count = count;
			m_reg_table = reg_table;
			break;
		case f64:
			m_fd_count = count;
			m_reg_table_64 = reg_table;
			break;
		default:
			break;
	}
	
	return  return_flag;
}

/*
 * Function Introduce: Another overloaded version of SetRegs().
 * 					   When gvinst's source operands is more than one, this version will be chosen.
 * 
 * Input Parameters: GVInst gvisnt 
 * Output Parameters: string dreg, vector<string> sreg
 * 		dreg is one of GPU's registers which is allocated to gvinst's destination operand. 
 * 		sreg corresponds to source opedrands.  
 * Returned Value: Always return 1.
 */
int PTXTranslator::SetRegs(const GVInst &gvinst, string &dreg, vector<string> &sreg)
{
	string reg, base;
	string soperand[2];
	int count, pos;
	map<string, string> reg_table;
	map<string, string>::iterator map_it;

	switch (gvinst.GetDType())
	{
		case s32:
			reg = "%r";
			count = m_r_count;
			reg_table = m_reg_table;
			break;
		case f32:
			reg = "%f";
			count = m_f_count;
			reg_table = m_reg_table;
			break;
		case f64:
			reg = "%fd";
			count = m_fd_count;
			reg_table = m_reg_table_64;
			break;
		default:
			break;
	}
	//Set sregs
	soperand[0] = gvinst.GetSoperand1();
	soperand[1] = gvinst.GetSoperand2();
	int snum = 2;
	while (snum--)
	{
		if (soperand[1 - snum].substr(0, 2) == "0x")
			sreg.push_back(soperand[1 - snum]);
		else if (soperand[1 - snum] == "v0")
			sreg.push_back("0x0");
		else if (soperand[1 - snum] == "f0")
			sreg.push_back("0f00000000");
		else 
		{	
			map_it = reg_table.find(soperand[1-snum]);
			if (map_it != reg_table.end())
				sreg.push_back(map_it->second);
			else
			{
				sreg.push_back(reg + int2str(count));
				reg_table.insert(make_pair(soperand[1-snum], sreg.back()));
				++count;
			}
		}
	}
	//Set dreg
	base = gvinst.GetDoperand();
	pos= base.find('(');
	if (pos != string::npos)
		base = base.substr(0, pos);
	
	map_it = reg_table.find(base);
	if (map_it != reg_table.end())
		dreg = map_it->second;
	else
	{
		dreg = reg + int2str(count);
		reg_table.insert(make_pair(base, dreg));
		count++;
	}

	switch (gvinst.GetDType())
	{
		case s32:
			m_r_count = count;
			m_reg_table = reg_table;
			break;
		case f32:
			m_f_count = count;
			m_reg_table = reg_table;
			break;
		case f64:
			m_fd_count = count;
			m_reg_table_64 = reg_table;
			break;
		default:
			break;
	}
	
	return 1;
}

/*
 * Function Introduce: Map labels to the basic blocks' entry addresses.
 *
 * Input Parameters: GVInst gvinst
 * Output Parameters: string loop_label
 *		loop_label is a string map to a basic block's entry address, format with "$loop_X".
 *		X stands for the number of loop labels had been mapped.
 * Returned Value: Always return 1.
 */
int PTXTranslator::SetLoops(const GVInst &gvinst, string &loop_label)
{
	string offset;
	string loop_prefix = "$loop_";
	int bpos, epos, count;
	offset = gvinst.GetDoperand();
	bpos = offset.find('(');
	epos = offset.find(')');
	if (bpos != string::npos && epos != string::npos)
		offset = offset.substr(bpos+1, epos-bpos-1);

	map<string, string>::iterator map_it;
	map<string, string> reg_table = m_addr_loop_label;
	count = m_loop_count;
	map_it = reg_table.find(offset);
	if (map_it != reg_table.end())
		loop_label = map_it->second;
	else
	{
		loop_label = loop_prefix + int2str(count);
		reg_table.insert(make_pair(offset, loop_label));
		count++;
	}
	m_addr_loop_label = reg_table;
	m_loop_count = count;
	
	return 1;
}

/*
 * Function Introduce: Encode GVInsts' MOV, GET adn PUT instructions.
 *                     Output PTX codes to file assigned in class PTXTranslator.
 * Encode Sample: MOV.s32	vi, vj => mov.s32	%ry, %rx;
 * 				  Note: %rx is GPU's register allocate to vi;
 * 				  		%ry is GPU's register allocate to vj.
 * Input Parameters: GVInst gvinst
 */
void PTXTranslator::EncodeMOV(const GVInst  &gvinst)
{
	// Do not encode PUT.f64
	if (gvinst.GetOpcode() == OP_PUT 
		&& gvinst.GetDType() == f64)
		return;

	string sreg, dreg;
	map<string, string>::iterator map_it;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);

	int flag = SetRegs(gvinst, dreg, sreg);
	switch (gvinst.GetDType())
	{
		case s32:
			if (flag == 2)
			{
#ifndef	FLOAT
				output << "cvt.rn.f32.s32	" << dreg << ", " << sreg << ";\n";
#endif
#ifdef FLOAT
				output << "mov.b32	" << "%b0, " << sreg << ";\n";
				output << "mov.b32	" << dreg << ", %b0;\n";
			
#endif	
				break;
			}
			else if (flag == 3)
			{
#ifndef FLOAT
				output << "cvt.rni.s32.f32	" << dreg << ", " << sreg << ";\n";
#endif	

#ifdef	FLOAT
				output << "mov.b32	%b0, " << sreg << ";\n";
				output << "mov.b32	" << dreg << ", " << sreg << ";\n";
				
#endif		

				break;
			}
			output << "mov.s32" << "	" << dreg<< ", " << sreg << ";\n";
			break;
		case f32:
			output << "mov.f32" << "	" << dreg<< ", " << sreg << "; \n";
			break;
		case f64:
			output << "mov.f64" << "	"
				   << dreg << ", " << sreg
				   << ";\n";
			break;
		default:
			break;
	}

	output.close();
}

/*
 * Function Introduce: Encode GVInsts' LI instruction.
 *                     Output PTX codes to file assigned in class PTXTranslator.
 * Encode Sample: LI.s32 0x1, vi => mov.s32 %rx, 0x1;
 * 				  Note: %rx is GPU's register allocate to vi.
 * Input Parameters: GVInst gvinst
 */
void PTXTranslator::EncodeLI(const GVInst &gvinst)
{
	string sreg, dreg;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);

	SetRegs(gvinst, dreg, sreg);
	output << "mov.s32" << "	"
		   << dreg << ", "
		   << sreg << "; \n";
	output.close();
}

/*
 * Function Introduce: Encode GVInsts' LD instruction.
 *                     Output PTX codes to file assigned in class PTXTranslator.
 * Encode Sample: LD.s32	vi, vj => mov.s32	%ry, %rx;
 * 				  Note: %rx is GPU's register allocate to vi;
 * 				  		%ry is GPU's register allocate to vj.
 * 				  LD.s32	vi(0x0), fj => ld.global.f32	%f*, [%rx+0];
 * 				  						   cvt.rni.s32.f32	%ry, %f*;
 * 				  Note: %f* is GPU.s register allocate to ftemp.
 * 				  LD.f32	vi(0x0), fj => ld.global.f32	%fy, [%rx+0];
 * Input Parameters: GVInst gvinst
 */
void PTXTranslator::EncodeLD(const GVInst &gvinst)
{
	string sreg, dreg, freg;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);
	
	int flag = SetRegs(gvinst, dreg, sreg);
	if (gvinst.GetDType() == s32)
	{
		if (!flag)
		{
		#ifndef	FLOAT
			map<string, string>::iterator map_it = m_reg_table.find("ftemp");
			if (map_it != m_reg_table.end())
				freg = map_it->second;
			else
			{
				freg = "%f" + int2str(m_f_count);
				m_reg_table.insert(make_pair("ftemp", freg));
				m_f_count++;
			}
			output << "ld.global.f32" 
				   << "	" << freg 
				   << ", " << "[" << sreg 
				   << "]" << ";\n"; 
			output << "cvt.rni.s32.f32" << " " << dreg << ", "
				   << freg << ";\n";
		#endif
		#ifdef FLOAT
			output << "ld.global.b32	%b0, [" 
				   << sreg << "]" << ";\n"; 
			output << "mov.b32	" << dreg << ", %b0;\n";
		#endif
		}
		else
		{
			output << "mov.s32" << "	"
			 	   << dreg << ", " 
				   << sreg << ";\n";	
		}
	}
	else if (gvinst.GetDType() == f32)
	{
		if (!flag)
		{
			output << "ld.global.f32" 
				   << "	" << dreg 
				   << ", " << "[" << sreg 
				   << "]" << ";\n"; 
		}
		else
		{
			output << "mov.f32" << "	"
			       << dreg << ", "
				   << sreg << ";\n";
		}
	}

	output.close();
}

/*
 * Function Introduce: Encode GVInsts' ST instruction.
 *                     Output PTX codes to file assigned in class PTXTranslator.
 * Encode Sample: ST.s32	vi, vj(0x0) 
 * 				  => cvt.rn.f32.s32		%fz, %rx;
 * 				     st.global.f32		[%ry+0], %fz;
 * 				  Note: %rx is GPU's register allocate to vi;
 * 				  		%ry is GPU's register allocate to vj;
 * 				  		%fz is GPU's floating point register.
 * 				  ST.f32	fi, vj(0x0) => st.global.f32	[%ry+0], %fx;
 * Input Parameters: GVInst gvinst
 */
void PTXTranslator::EncodeST(const GVInst &gvinst)
{
	string sreg, dreg, freg;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);
	
	int flag = SetRegs(gvinst, dreg, sreg);
	if (gvinst.GetDType() == s32)
	{
		if (!flag)
		{
#ifndef FLOAT
			map<string, string>::iterator map_it = m_reg_table.find("ftemp");
			if (map_it != m_reg_table.end())
				freg = map_it->second;
			else
			{
				freg = "%f" + int2str(m_f_count);
				m_reg_table.insert(make_pair("ftemp", freg));
				m_f_count++;
			}
			output << "cvt.rn.f32.s32" << "	" <<  freg << ", " << sreg
				   << ";\n";
			output << "st.global.f32" << "	[" << dreg << "], " << freg
				   << ";\n";
#endif
#ifdef	FLOAT
			output << "mov.b32	%b0, " << sreg << ";\n";
			output << "st.global.b32 [" << dreg << "], %b0;\n";
#endif
		}
		else
		{
			output << "mov.s32" << "	"
				   << dreg << ", "
				   << sreg << ";\n";
		}
	}
	else if (gvinst.GetDType() == f32)
	{
		if (!flag)
		{
			output << "st.global.f32" << "	"
				   << "[" << dreg << "], " 
				   << sreg << ";\n"; 
		}
		else
		{
			output << "mov.f32" << "	"
			       << dreg << ", "
				   << sreg << ";\n";
		}
	}

	output.close();
}

/*
 * Function Introduce: Encode GVInsts' ADD instruction.
 *                     Output PTX codes to file assigned in class PTXTranslator.
 * Encode Sample: ADD.s32	vi, vj, vk => add.s32	%rz, %rx, %ry;
 * 				  Note: %rx is GPU's register allocate to vi;
 * 				  		%ry is GPU's register allocate to vj;
 * 				  		%rz is GPU's register allocate to vk.
 * Input Parameters: GVInst gvinst
 */
void PTXTranslator::EncodeADD(const GVInst &gvinst)
{
	string dreg;
	vector<string> sreg;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);

	SetRegs(gvinst, dreg, sreg);
	if (gvinst.GetDType() == s32)
		output << "add.s32";
	else if (gvinst.GetDType() == f32)
		output << "add.f32";
	else if (gvinst.GetDType() == f64)
		output << "add.f64";
	output << "	" << dreg << ", "
		   << sreg[0] << ", " << sreg[1] << ";\n";
	output.close();
}

/*
 * Function Introduce: Encode GVInsts' ADDI instruction.
 *                     Output PTX codes to file assigned in class PTXTranslator.
 * Encode Sample: ADDI.s32	vi, 0x0, vj => add.s32	%ry, %rx, 0x0;
 * 				  Note: %rx is GPU's register allocate to vi;
 * 				  		%ry is GPU's register allocate to vj.
 * Input Parameters: GVInst gvinst
 */
void PTXTranslator::EncodeADDI(const GVInst &gvinst)
{
	string dreg;
	vector<string> sreg;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);

	SetRegs(gvinst, dreg, sreg);
	output << "add.s32" << "	"
		   << dreg << ", "
		   << sreg[0] << ", " << sreg[1] << ";\n";
	
	output.close();
}

/*
 * Function Introduce: Encode GVInsts' SUB instruction.
 *                     Output PTX codes to file assigned in class PTXTranslator.
 * Encode Sample: SUB.s32	vi, vj, vk => sub.s32	%rz, %rx, %ry;
 * 				  Note: %rx is GPU's register allocate to vi;
 * 				  		%ry is GPU's register allocate to vj;
 * 				  		%rz is GPU's register allocate to vk.
 * Input Parameters: GVInst gvinst
 */
void PTXTranslator::EncodeSUB(const GVInst &gvinst)
{
	string dreg;
	vector<string> sreg;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);

	SetRegs(gvinst, dreg, sreg);
	if (gvinst.GetDType() == s32)
		output << "sub.s32";
	else if (gvinst.GetDType() == f32)
		output << "sub.f32";
	else if (gvinst.GetDType() == f64)
		output << "sub.f64";
	output << "	"
		   << dreg << ", "
		   << sreg[0] << ", " 
		   << sreg[1] << ";\n";

	output.close();
}

/*
 * Function Introduce: Encode GVInsts' SUBI instruction.
 *                     Output PTX codes to file assigned in class PTXTranslator.
 * Encode Sample: SUBI.s32	vi, 0x0, vj => sub.s32	%ry, %rx, 0x0;
 * 				  Note: %rx is GPU's register allocate to vi;
 * 				  		%ry is GPU's register allocate to vj;
 * Input Parameters: GVInst gvinst
 */
void PTXTranslator::EncodeSUBI(const GVInst &gvinst)
{
	string dreg;
	vector<string> sreg;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);

	SetRegs(gvinst, dreg, sreg);
	output << "sub.s32" << "	"
		   << dreg << ", "
		   << sreg[0] << ", " 
		   << sreg[1] << ";\n";

	output.close();
}

/*
 * Function Introduce: Encode GVInsts' AND instruction.
 *                     Output PTX codes to file assigned in class PTXTranslator.
 * Encode Sample: AND.s32	vi, vj, vk => and.b32	%rz, %rx, %ry;
 * 				  Note: %rx is GPU's register allocate to vi;
 * 				  		%ry is GPU's register allocate to vj;
 * 				  		%rz is GPU's register allocate to vk.
 * Input Parameters: GVInst gvinst
 */
void PTXTranslator::EncodeAND(const GVInst &gvinst)
{
	string dreg;
	vector<string> sreg;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);

	SetRegs(gvinst, dreg, sreg);
	output << "and.b32" << "	"
		   << dreg << ", "
		   << sreg[0] << ", "
		   << sreg[1] << ";\n";

	output.close();
}

/*
 * Function Introduce: Encode GVInsts' MUL instruction.
 *                     Output PTX codes to file assigned in class PTXTranslator.
 * Encode Sample: MUL.s32	vi, vj, vk => mul.lo.s32	%rz, %rx, %ry;
 * 				  Note: %rx is GPU's register allocate to vi;
 * 				  		%ry is GPU's register allocate to vj;
 * 				  		%rz is GPU's register allocate to vk.
 * 				  MUL.f64	fi, fj, NA, fk, NA => mul.f64	%fdz, %fdx, %fdy;
 * 				  Note: %fdx mapped to fi;
 * 				  		%fdy mapped to fj;
 * 				  		%fdz mapped to fz.
 * Input Parameters: GVInst gvinst
 */
void PTXTranslator::EncodeMUL(const GVInst &gvinst)
{
	string dreg;
	vector<string> sreg;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);

	SetRegs(gvinst, dreg, sreg);
	switch (gvinst.GetDType())
	{
		case s32:
			output << "mul.lo.s32";
			break;
		case f32:
			output << "mul.lo.f32";
			break;
		case f64:
			output << "mul.f64";
			break;
		default:
			break;
	}
	output << "	" << dreg << ", "
		   << sreg[0] << ", " << sreg[1] << ";\n";

	output.close();
}

/*
 * Function Introduce: Encode GVInsts' SLL instruction.
 *                     Output PTX codes to file assigned in class PTXTranslator.
 * Encode Sample: SLL.s32	vi, vj, vk => shb.b32	%rz, %rx, %ry;
 * 				  Note: %rx is GPU's register allocate to vi;
 * 				  		%ry is GPU's register allocate to vj;
 * 				  		%rz is GPU's register allocate to vk.
 * Input Parameters: GVInst gvinst
 */
void PTXTranslator::EncodeSLL(const GVInst &gvinst)
{
	string dreg;
	vector<string> sreg;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);

	SetRegs(gvinst, dreg, sreg);
	output.seekp(ios::end);

	output << "shl.b32" << "	"
		   << dreg << ", "
		   << sreg[0] << ", " <<  sreg[1] << "; \n";
	output.close();
}

/*
 * Function Introduce: Encode GVInsts' CONV instruction.
 *                     Output PTX codes to file assigned in class PTXTranslator.
 * Encode Sample: CONV.f32.f64	fi, fj => cvt.f64.f32	%fdy, %fx;
 * 				  Note: %fx is GPU's register allocate to fi;
 * 				  		%fdy is GPU's register allocate to fj.
 * 				  CONV.f64.f32	fi, fj => cvt.rn.f32.f64	%fy, %fdx;
 * 				  Note: %fdx is GPU's register allocate to fi;
 * 				  		%fy is GPU's register allocate to fj.
 * Input Parameters: GVInst gvinst
 */
void PTXTranslator::EncodeCONV(const GVInst &gvinst)
{
	string dreg, sreg;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);
	DataType soperand_type, doperand_type;
	soperand_type = gvinst.GetSType();
	doperand_type = gvinst.GetDType();
	//map<string, string>::iterator map_it;

	SetRegs(gvinst, dreg, sreg);
	if (doperand_type == f32 
	    && soperand_type == f64)
		output << "cvt.rn";
	else
		output << "cvt";
	switch (doperand_type)
	{
		case f32:
			output << ".f32";
			break;
		case f64:
			output << ".f64";
			break;
		default:
			break;
	}
	switch (soperand_type)
	{
		case f32:
			output << ".f32";
			break;
		case f64:
			output << ".f64";
			break;
		default:
			break;
	}
	output << "	" << dreg << ", "
		   << sreg << ";\n";

	output.close();
}

/*
 * Function Introduce: Encode GVInsts' CMP instruction.
 *                     Output PTX codes to file assigned in class PTXTranslator.
 * Encode Sample: CMP cc vi, vj, vk => set.cc.s32.s32	%rz, %rx, %ry;
 * 				  Note: cc is CMP instruction's condition code.
 * 				  		%rx is GPU's register allocate to vi;
 * 				  		%ry is GPU's register allocate to vj;
 * 				  		%rz is GPU's register allocate to vk.
 * Input Parameters: GVInst gvinst
 */
void PTXTranslator::EncodeCMP(const GVInst &gvinst)
{
	string dreg;
	vector<string> sreg;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);

	SetRegs(gvinst, dreg, sreg);
	CC condition_code = gvinst.GetCC();
	switch(condition_code)
	{
		case eq:
			output << "set.eq.s32.s32" << "	";
			break;
		case ne:
			output << "set.ne.s32.s32" << "	";
			break;
		case lt_u32:// BE
			output << "set.lt.u32.u32" << "	";
			break;
		case lt_s32:
			output << "set.lt.s32.s32" << "	";
			break;
		case le:
			output << "set.le.s32.s32" << "	";
			break;
		case gt:
			output << "set.gt.s32.s32" << "	";
			break;	
		case ge:
			output << "set.ge.s32.s32" << "	";
			break;
		default:
			break;
	}

	output << dreg << ", " 
		   << sreg[0] << ", "
		   << sreg[1] << "; \n";
	// If result is True, 0xFFFFFFFF is written
	output << "and.b32" << "	"
		   << dreg << ", "
		   << dreg << ", " << "0x00000001;\n"; 
	
	output.close();
}

/*
 * Function Introduce: Encode GVInsts' BRANCH instruction.
 *                     Output PTX codes to file assigned in class PTXTranslator.
 * Encode Sample: BRANCH cc vi, vj, vk(0xXXXXXXXX) 
 * 				  => setp.cc.s32	%p, %rx, %ry;
 * 				     @%p bra	$loop_X;
 * 				  Note: cc is CMP instruction's condition code.
 * 				  		$loop_X mapped to [vk + 0xXXXXXXXX] address
 * 				  		%rx is GPU's register allocate to vi;
 * 				  		%ry is GPU's register allocate to vj.
 * Input Parameters: GVInst gvinst
 */
void PTXTranslator::EncodeBRANCH(const GVInst &gvinst)
{
	string dreg;
	vector<string> sreg;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);
	string loop_label;

	SetRegs(gvinst, dreg, sreg);
	SetLoops(gvinst, loop_label);
	switch (gvinst.GetCC())
	{
		case eq:
			output << "setp.eq.s32" << "	";
			break;
		case ne:
			output << "setp.ne.s32" << "	";
			break;
		case lt_u32:
			output << "setp.lt.u32" << "	";
			break;
		case lt_s32:
			output << "setp.lt.s32" << "	";
			break;
		case le:
			output << "setp.le.s32" << "	";
			break;
		default:
			break;
	}

	output << "%p" << m_p_count << ", "
		   << sreg[0] << ", " << sreg[1]
		   << ";\n";
	output << "@%p" << m_p_count << " bra	"
		   << loop_label << ";\n";
	++m_p_count;
	
	output.close();
}

/*
 * Function Introduce: Encode the string line begins with "entry:".
 *                     Output PTX codes to file assigned in class PTXTranslator.
 * Encode Sample: entry: 0xXXXXXXXX => $loop_X:
 * 				  Note: $loop_X mapped to [0xXXXXXXXX] address.
 * 				  		"entry:" isn't included in GVInst, but it is necessary to label the basic block's entry address.  
 * Input Parameters: GVInst gvinst
 */
void PTXTranslator::EncodeENTRY(const GVInst &gvinst)
{
	string loop_label;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);
	SetLoops(gvinst, loop_label);
	output << loop_label << ":\n";
	
	//Edited by Dylan 10.4.8
	//add vector "m_entry_loop"
	vector<string>::iterator vec_it;
	vector<string> vec_table = m_entry_loop;
	//count = m_loop_count;
	vec_it = find(vec_table.begin(), vec_table.end(), loop_label);

	if (vec_it == vec_table.end())
	{
		//loop_label = loop_prefix + int2str(count);
		vec_table.push_back(loop_label);
	}
	m_entry_loop = vec_table;
	
	output.close();
}

/*
 * Function Introduce: Encode the string line begins with "next:".
 *                     Output PTX codes to file assigned in class PTXTranslator.
 * Encode Sample: next: 0xXXXXXXXX => bra.uni	$loop_X;
 * 				  Note: $loop_X mapped to [0xXXXXXXXX] address.
 * 				  		"next:" isn't included in GVInst, but it is necessary to label the next basic block's entry address which will be jumpped to.  
 * Input Parameters: GVInst gvinst
 */
void PTXTranslator::EncodeNEXT(const GVInst &gvinst)
{
	string loop_label;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);
	SetLoops(gvinst, loop_label);
	output << "bra.uni" << "	" << loop_label << ";\n\n\n";

	output.close();
}

/*
 * Function Introduce: Called when  no corresponding encode function.
 *                     Output PTX comments to file assigned in class PTXTranslator.
 * Encode Sample: # No corresponding encode function!;
 */
void PTXTranslator::EncodeERR()
{
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);
	output << "// No corresponding encode function!\n";

	output.close();
}

/*
 * Function Introduce: Traverse map container and output its' info to assigned file.
 * 					   Here we use the function to output the relations between GPU's registers and gvinsts' operands
 * Input Parameters: string t_file, map<string, string> table
 * 					 "t_file" is the assigned file's filename.
 * 					 "table" is the map container need be traversed.
 */
void PTXTranslator::MapTraversal(const string &t_file, const map<string, string> &table)
{
	map<string, string>::const_iterator map_it = table.begin();
	ofstream output(t_file.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);

	while (map_it != table.end())
	{
		string blanks(10 - (map_it->first).length(), ' ');
		output << map_it->first << blanks << "=>  "
			   << map_it->second << endl;
		++map_it;
	}
	output << endl;

	output.close();
}

/*
 * Function Introduce: Initial an object of class GVInst named gvisnt, and choose the encode function according to gvinst's opcode.
 * Input Parameters: list<string> line_of_gvinst
 * 					 "line_of_gvinst" contains a GVInst instruction, used for initialize an object of class GVInst.
 */
void PTXTranslator::Encode(const list<string> &line_of_gvinst)
{
	GVInst gvinst(line_of_gvinst);
	Opcode opcode = gvinst.GetOpcode();
	
	switch (opcode)
	{
		case OP_GET:
		case OP_PUT:
		case OP_MOV:
			EncodeMOV(gvinst);
			break;
		case OP_LD:
			EncodeLD(gvinst);
			break;
		case OP_ST:
			EncodeST(gvinst);
			break;
		case OP_LI:
			EncodeLI(gvinst);
			break;
		case OP_ADD:
			EncodeADD(gvinst);
			break;
		case OP_ADDI:
			EncodeADDI(gvinst);
			break;
		case OP_SUB:
			EncodeSUB(gvinst);
			break;
		case OP_SUBI:
			EncodeSUBI(gvinst);
			break;
		case OP_MUL:
			EncodeMUL(gvinst);
			break;
		case OP_SLL:
			EncodeSLL(gvinst);
			break;
		case OP_AND:
			EncodeAND(gvinst);
			break;
		case OP_CONV:
			EncodeCONV(gvinst);
			break;
		case OP_CMP:
			EncodeCMP(gvinst);
			break;
		case OP_BRANCH:
			EncodeBRANCH(gvinst);
			break;
		case OP_ENTRY: 
			EncodeENTRY(gvinst);
			break;
		case OP_NEXT: 
			EncodeNEXT(gvinst);
			break;
		case OP_ERR:
			//cout << line_of_gvinst.front() << endl;
			//assert(0);
			EncodeERR();
			break;
		default:
			break;
	}
}

/*
 * Function Introduce: Start translation from GVInst to PTX.
 * Input Parameters: Record *record_ptr, string infile, string outfile
 * 					 "record_ptr" will be used when call PTXInit();
 * 					 "infile" is GVInst file's filename.
 * Output Parameters: "outfile"
 * 					  The file nameed "outfile" stores the created PTX codes.
 * Returned Value: Always return 1.
 */
void PTXTranslator::TransEngine(Record *record_ptr, const string &infile, const string &outfile)
{
	list<string> line_of_gvinst;
	string line, words, op;
	string numerics("0123456789");
	int current_bb = -1;
	int op_pos, last_st_pos, line_num = 0;

	ifstream input(infile.c_str(), ios::binary);
	// Edit by Dylan
	// Data: 2010/3/13
	//search the last ST
 	//cout << "trans engine start ! " << endl;	
	while (getline(input, line))
	{
		line_num++;
		if (line.length())
		{
			string::size_type pos = line.find("ST.");
			if (pos != string::npos)
				last_st_pos = line_num;
		}
	}
	input.clear();
	input.seekg(0, ios::beg);

	//flash the destination file
	ofstream output(m_tempfile.c_str(), ios::binary | ios::trunc);
	output.close();

	output.open(m_tempfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);

	PTXInit(record_ptr);
	
	while (getline(input, line))
	{
		output.seekp(ios::end);
		if (line.length())
		{
			// Erase the first Basic Block =>for(j = 0;...)
			op_pos = line.find_first_of(" ");
			op = line.substr(0, op_pos);
			if (op == "entry:")
			{
				current_bb++;
				if (!current_bb)
					output << "\n// The first Basic Block had been erased.\n" << endl;
			}
			/*
			if(op == "next:")
			{
				if(string::npos != line.find("loop_8"))
				{
					m_entry_pos = output.tellp();
						output << "dylan" << m_entry_pos << endl;
				}
			}*/
			if (current_bb)
			{
				int global_var = 0;
				int count = 6;
				if (line == "entry: 0x8048273")
				{
					global_var = 1;
				}
				output << "// * " << line << endl;  
				replace_all(line, ",", " ");
				replace_all(line, ".", " ");
				replace_all(line, "NA", " ");
				// Load GVInst info(opcode & oprands) to List line_of_gvinst
				istringstream istr(line);
				while (istr >> words)
					line_of_gvinst.push_back(words);
				// Modify List of gvinst for removing line index
				if (line_of_gvinst.front().find_first_of(numerics) != string::npos)
					line_of_gvinst.pop_front();
				// Send to encode into PTX
				Encode(line_of_gvinst);
				
				while (global_var)
				{
					getline(input, line);
					if(!count--)
					{
						line_of_gvinst.clear();
						line_of_gvinst.push_back("MOV");
						line_of_gvinst.push_back("s32");
						line_of_gvinst.push_back("0x0");
						line_of_gvinst.push_back("v1");
						Encode(line_of_gvinst);
						line_of_gvinst.clear();
						line_of_gvinst.push_back("MOV");
						line_of_gvinst.push_back("s32");
						line_of_gvinst.push_back("0x0");
						line_of_gvinst.push_back("s2");
						Encode(line_of_gvinst);
						break;
					}
				}
				// Edit by Dylan
				// Date: 2010/3/13
				/*
				list<string>::iterator list_it = line_of_gvinst.begin(); 
				if (*(list_it++) == "ST")
				{
					output << "// GVInsts behind last ST is needless.\n" << endl;
					break;
				}
				*/
			}
			// Delete GVInsts behind last ST
			
			if (!--last_st_pos)
			{
				output << "// GVInsts behind last ST is needless.\n" << endl;
				output << "/* GVInst Objects num: " << GVInst::GVInstNum() << " */\n" << endl;
				break;
			}


		}
		line_of_gvinst.clear();
	}
	// Check if each loop has entry
			
	if (m_loop_count > m_entry_loop.size())
	{
		map<string, string>::const_iterator map_it = m_addr_loop_label.begin();
		while (map_it != m_addr_loop_label.end())
		{
			vector<string>::iterator iter = 
				find(m_entry_loop.begin(), m_entry_loop.end(), map_it->second);
			if (iter != m_entry_loop.end())
			{
				//cout << *iter << m_entry_loop.size() << endl;
				m_entry_loop.erase(iter);
			}
			else
			{
				m_entry_loop.push_back(map_it->second);
				//cout << map_it->second;
			}
			map_it++;
		}
		//cout << m_entry_loop.back() << m_entry_loop.size()<< endl;
	}	
	PTXOut(outfile);
	output << endl;

	input.close();
	output.close();
}

/*
 * Function Introduce: Before tranlation, output PTX's initial info like .version, .target and PTX function's initial info, including function name, formal parameter, mapping loop index, load parameters, etc.
 * Input Parameters: Record *record_ptr
 * 					 Piont to an array records PTX function's parameters
 * Returned Value: Always return 1.
 */
int PTXTranslator::PTXInit(Record *record_ptr)
{
	int records_num = Record::record_num;
	string param_name;
	ofstream output(m_tempfile.c_str(), ios::binary | ios::app);
	
	output << ".version 1.4\n"
		   << ".target sm_10, map_f64_to_f32\n\n"
		   << ".entry MatrixMulKernel(\n"
		   << ".param .u32 regs,\n";

	//Edit by Dylan 10.1.20
	//Set Params of PTX kernel function
	for (int i = 0; i < records_num - 1; i++)
	{
		param_name = record_ptr[i].getLabel();
		output << ".param .u32 " << param_name;
		if (i == records_num - 2)
			output << "){\n\n";
		else
			output << ",\n";
	}

	m_reg_init_pos = output.tellp();
	//Edit by Dylan 10.3.12
	//add ARG0, ARG1, TidA
	output << "cvt.u32.u16	%r0, %tid.x; //threadIdx.x\n"
		   << "cvt.u32.u16	%r1, %tid.y; //threadIdx.y\n"
		   << "mov.u16		%rh0, %ntid.x; //blockDim.x\n"
		   << "mov.u16		%rh1, %ctaid.x; //blockIdx.x\n"
		   << "mov.u16		%rh2, %ntid.y; //blockDim.y\n"
		   << "mov.u16		%rh3, %ctaid.y; //blockIdx.y\n"
		   << "mul.wide.u16	%r4, %rh0, %rh1;\n"
		   //<< "add.u32		%r2, %r4, %r0; //j=blockDim.x*blockIdx.x+threadidx.x\n"
		   << "add.u32		%r2, %r4, %r0; //x=blockDim.x*blockIdx.x+threadidx.x\n"
		   <<"mul.wide.u16	%r4, %rh2, %rh3;\n"
		   //<<"add.u32		%r3, %r4, %r1; //i=blockDim.y*blockIdx.y+threadIdx.y\n"
		   <<"add.u32		%r3, %r4, %r1; //y=blockDim.y*blockIdx.y+threadIdx.y\n"
		   << "mov.s32		%r4, 0x0; //v0 = %r4 = 0\n";
	m_reg_table.insert(make_pair("%tid.x", "%r0"));
	m_reg_table.insert(make_pair("%tid.y", "%r1"));
	m_reg_table.insert(make_pair("TidX", "%r3"));
	m_reg_table.insert(make_pair("TidY", "%r2"));
	m_reg_table.insert(make_pair("v0", "%r4"));
	// insert ftemp
	m_reg_table.insert(make_pair("ftemp", "%f0"));
	m_reg_table.insert(make_pair("rtemp", "%r5"));
	
	m_r_count = 6;

	// Load params should been loaded
	for (int i = 0; i < records_num - 1; i++)
	{
		param_name = record_ptr[i].getLabel();
		if (param_name.substr(0,4) == "TEMP")
			continue;
		output << "ld.param.u32	%r" 
			   << m_r_count << ", ["
			   << param_name << "];\n";
	//Edit by Dylan 10.3.12
		if (param_name == "WA" || param_name == "WB" 
			|| param_name == "HA" || param_name == "HB" 
			|| param_name == "ARG0" || param_name == "ARG1")
			output << "ld.global.s32	%r" << m_r_count << ", [%r" << m_r_count << "];\n";
		m_reg_table.insert(make_pair(param_name, "%r" + int2str(m_r_count)));
		m_r_count++;
	}	
	m_rh_count = 4;
	//m_f_count = 0;
	m_f_count = 1;
	m_fd_count = 0;
	
	output.close();
	return 1;
}

/*
 * Function Introduce: Output ptx codes for pre-allocating GPU's registers.
 * 					   Called after translation, since the registers number need to be allocted could be decidied when the translation is finished.
 * Output Parameters: string outfile
 * 					  The file nameed "outfile" stores the created PTX codes.
 * Returned Value: Always return 1.
 */
int PTXTranslator::RegInit(const string &outfile)
{
	ofstream output(outfile.c_str(), ios::binary | ios::app);
	output.seekp(ios::end);
	output << ".reg .u16 %rh<" << m_rh_count << ">;\n"
		   << ".reg .u32 %r<" << m_r_count << ">;\n"
		   << ".reg .f32 %f<" << m_f_count << ">;\n"
		   << ".reg .f64 %fd<" << m_fd_count << ">;\n"
		   << ".reg .pred %p<" << m_p_count << ">;\n"
		   << ".reg	.b32 %b<1>;\n"
		   << endl;

	output.close();
	return 1;
}

/*
 * Function Introduce: Output the whole ptx codes to the assigned file.
 * Output Parameters: string outfile
 * 					  The file nameed "outfile" stores the created PTX codes.
 * Returned Value: Always return 1.
 */
int PTXTranslator::PTXOut(const string &outfile)
{
	//ifstream in_file(m_tempfile.c_str(), ios::binary);
	ifstream in_file;
	//flash the destination file
	ofstream out_file(outfile.c_str(), ios::binary | ios::trunc);
	out_file.close();
	out_file.open(outfile.c_str(), ios::binary | ios::app);
	char ch;
	string line, loop;
	int loop_num = m_entry_loop.size();
	cout << loop_num << endl;
	
	for (int i = 0; i < loop_num; i++)
	{
	in_file.open(m_tempfile.c_str(), ios::binary);
	loop = m_entry_loop[i];
	string::size_type pos;
	int line_num = 0;
	while (getline(in_file, line))
	{
		pos = line.find(loop);
		if (pos != string::npos)
			line_num = in_file.tellg();
	}
	if (line_num != 0)
		m_entry_pos.push_back(line_num);
	in_file.close();
	}
	//cout << m_entry_pos[0] << endl;
	in_file.open(m_tempfile.c_str(), ios::binary);
	while (in_file.get(ch))
	{
		int i = in_file.tellg();
		if (i == m_reg_init_pos)
		{
			out_file << endl;
			RegInit(outfile);
			out_file << "/*********************************\n";
			out_file << ">>>>>>>>>>>>>>>REG_MAP<<<<<<<<<<<<\n";
			out_file << endl;
			out_file << "<<64bits registers>>" << endl;
			MapTraversal(outfile, m_reg_table_64);
			out_file << "<<32bits registers>>" << endl;
			MapTraversal(outfile, m_reg_table);
			out_file << "<<loop labels>>" << endl;
			MapTraversal(outfile, m_addr_loop_label);
			out_file << "**********************************/\n";
		}
		
		for (int j = 0; j < loop_num; j++)
		if (i == m_entry_pos[j])
		{
			out_file << endl;
			//out_file << ch << "donguoxing " << m_entry_loop << " " << m_entry_loop.size() << " " << m_entry_pos << ":\n" << endl;
			//for(int j = 0; j < m_entry_loop.size(); j++)
			out_file << m_entry_loop[j] << ":" << endl;

		}
		out_file << ch;
	}

	out_file << "exit;\n" << "}\n" << endl;

	in_file.close();
	out_file.close();
	//system("cp " + outfile + " ptx.temp");
	system("rm -f ptx.temp");
	//system("cp  ./script/PTX/* ./temp");
	return 1;
}

